Dynamic Batching 分析 | Knowledge Wiki

1. Dynamic Batching

为什么动态凑批（Dynamic Batching）能提升显存带宽利用率？

在 GPU 推理中，每一层的计算都涉及两个动作：加载数据（显存 → 寄存器）和执行计算（ALU/Tensor Core）。

权重复用： 对于小模型，权重（Weights）通常比输入数据（Activations）大得多。在 $BS=1$ 时，GPU 每读入一次权重，只为一个请求服务。
摊薄开销： 动态凑批的核心逻辑是：“一次读取，多次使用”。
- 当 $BS=N$ 时，GPU 将模型权重从显存读取到片上缓存（L1/SRAM）后，可以连续为 $N$ 个输入请求进行计算。
- 显存带宽消耗主要由读取权重的流量决定。既然读取权重的流量在 $BS=1$ 和 $BS=N$ 时几乎相同（对于不大的 $N$ ），那么 $BS$ 越大，每一单位显存带宽所支撑的有效计算量（吞吐量）就越高。

为什么小模型在 $BS=1$ 时是典型的 Memory-bound？

这可以用 Roofline Model 来解释。

计算性能受两个指标限制：

算力峰值 ( $\pi$ )：GPU 每秒能做多少次浮点运算（TFLOPS）。
带宽峰值 ( $\beta$ )：GPU 每秒能搬运多少字节数据（GB/s）。

对于一个算子，我们定义其计算密度（Arithmetic Intensity, $I$ ）：

$I = \frac{\text{FLOPs（计算量）}}{\text{Bytes（访存量）}}$

小模型的特征： 参数量少，计算层浅。在 $BS=1$ 时， $I$ 非常小。
瓶颈判定： 当 $I < \frac{\pi}{\beta}$ 时，GPU 处于 Memory-bound 区域。此时，计算单元大部分时间在“等数据吃”，空有强大的算力却无用武之地。
结果： 此时即便你换一个算力强 10 倍的 GPU，推理延迟可能几乎不变，因为瓶颈在搬运数据的速度（带宽）上。

如何观察单模型推理中，哪些算子是 Memory-bound？

你可以使用 NVIDIA 的专业工具链进行定性与定量分析：

NVIDIA Nsight Systems (Nsys)
- 观察时间轴。如果内核（Kernel）执行时间很短，但内核之间的间隙（Gap）很大，或者大量的 cudaMemcpy 占据了时间轴，说明调度或访存是瓶颈。
NVIDIA Nsight Compute (Ncu)：深度诊断（推荐）
- 查看 SOL (Speed of Light) 指标。
- Memory Chart： 如果 DRAM Throughput 的百分比远高于 SM Utilization，说明该算子是 Memory-bound。
- Roofline Chart： Ncu 会直接画出 Roofline 图，并标出当前算子的点。如果点落在斜率线上（左侧），则是带宽受限；如果落在平顶线上（顶部），则是算力受限。

多模型混跑，一次 Profiling 如何区分不同的模型？

在混跑场景下，成千上万个 Kernel 在同一个 GPU 上交错执行，直接看底层 API 就像看乱码。区分它们的“金钥匙”是 NVTX (NVIDIA Tools Extension)。

NVTX 标记（最强手段）： 你可以在代码中为不同的模型打上标签。

C++
```
// C++ 示例
nvtxRangePush("Model_A_Inference");
run_model_a();
nvtxRangePop();
```
在 Nsight Systems 的时间轴上，你会看到一行专门的 NVTX Row，清晰地标出这段时间内执行的所有 Kernel 属于哪个模型。
CUDA Stream（流）区分： 通常不同的模型会运行在不同的 CUDA Stream 中。在 Nsys 中，你可以展开 Streams 视图。通过观察不同的 Stream ID 及其对应的 Kernel 序列，可以反推模型。
进程 PID/进程名： 如果你采用的是多进程部署（或通过 MPS 分发），Nsys 会按 Process 维度对 Kernel 进行分组展示。

Profiling 解决 GPU Gap

对于该情况，需要重点关注两点：

量化 Gap： 在 Nsight Systems 或其他 profiling 中，统计一下 Stream 上的空闲时间占比。如果超过 20%，必须优化。
检查异步性： 确认 cudaMemcpy 是否全部换成了 cudaMemcpyAsync，并且是否使用了 Pinned Memory（锁页内存）。

对于第二点，在 CUDA 编程中，将 cudaMemcpy 升级为 cudaMemcpyAsync 并配合 Pinned Memory，是解决 Trace 中“大 Gap”、提升 GPU 吞吐量的最核心手段。

简单来说，此举的目的是为了实现 “计算”与“通信”的重叠 (Overlap)，让 GPU 真正做到“手脚不停”。

1. Pinned Memory (锁页内存)：解决“隐形搬运”损耗

背景： 普通的 CPU 内存（Pageable Memory）是受操作系统管理的，可能会被换出到磁盘或在物理地址上移动。 问题： GPU 无法直接访问物理地址不固定的内存。当你调用标准的 cudaMemcpy 时，CUDA 驱动其实在背后偷偷做了两件事：

在 CPU 上开辟一块隐藏的 Pinned Memory。
将你的数据从普通内存拷贝到这块隐藏的 Pinned Memory。
再通过 DMA（直接存储器存取）将数据从 Pinned Memory 搬运到 GPU。

此举的目的：

减少一次拷贝： 如果你直接申请 cudaMallocHost (Pinned Memory)，数据就直接驻留在物理内存中。GPU 的 DMA 引擎可以直接“吸”走数据。
更高的带宽： Pinned Memory 通常能获得比 Pageable Memory 更高的 PCIe 传输带宽。
异步的前提： 这是最关键的一点。 CUDA 驱动要求，只有当内存是 Pinned 时，cudaMemcpyAsync 才能真正实现异步。否则，它依然会退化成同步阻塞模式

`cudaMemcpyAsync`：打破“串行”枷锁

此举的目的：

非阻塞 CPU： cudaMemcpy 会让 CPU 在那儿干等，直到数据搬完。cudaMemcpyAsync 发起指令后立即返回，CPU 可以接着去处理下一个 Batch 的凑批逻辑。
流水线化 (Pipelining)： 只有使用异步拷贝，你才能在同一个或不同的 CUDA Stream 中，让 “数据搬运 (HtoD)” 和 “核函数计算 (Kernel)” 在时间轴上重叠。

组合技的效果：从“走走停停”到“多线程并行”

让我们对比一下你在 Trace 中看到的景象：

方案 A：同步拷贝 + 普通内存（你目前可能的状态）

CPU 发起拷贝（阻塞） $\rightarrow$ 等待 $\rightarrow$ 拷贝完成 $\rightarrow$ CPU 下发 Kernel $\rightarrow$ GPU 计算 $\rightarrow$ 计算完成 $\rightarrow$ CPU 取回结果。

结果： GPU 只有在 Kernel 执行时才动，其他时间都在等数据，Trace 上表现为 Kernel 两侧有巨大的 Gap。

方案 B：异步拷贝 + 锁页内存（优化后的状态）

利用多个 Stream，你可以实现：

当 Kernel 正在计算 Batch N 的时候；
HtoD 引擎已经在搬运 Batch N+1 的输入数据；
DtoH 引擎已经在写回 Batch N-1 的计算结果。

此举的最终目的：

消除你看到的 Gap，让 GPU 的 SM 利用率（Compute） 和 Copy 引擎利用率 同时保持在高位。

pined memory

警告

cudaMemcpy 是否全部换成了 cudaMemcpyAsync，并且是否使用了 Pinned Memory?
已检查，目前业务中存在的 memcpy 分别是：cudaMemcpyAsync_ptsz_v7000、cudaMemcpyAsync_v3020 和 cudaMemsetAsync_v3020 三种，占总时间的比例不高。

这是一个非常典型的 “异步陷阱”。即便你在 Trace 中看到了 cudaMemcpyAsync 的符号，如果底层的内存不是 Pinned Memory，这个 API 在执行时往往会退化成同步行为，或者在 CPU 侧引入额外的开销。

如何确认是否使用了 Pinned Memory？

你可以通过以下三种方式进行“排雷”：

方法 A：代码审计（最直接）

检查你的代码（或是推理框架源码）中分配输入/输出 Buffer 的位置：

普通内存（Pageable）： 使用 malloc(), new, std::vector, 或者 Go 的 make([]byte)。
锁页内存（Pinned）： 必须显式调用 cudaMallocHost(), cudaHostAlloc()。
已分配内存转换： 是否调用过 cudaHostRegister() 将现有的普通内存临时“锁住”。

方法 B：查看 Nsight Systems 的详细属性

在 Nsight Systems 的时间轴上：

点击那条 cudaMemcpyAsync 的色块。
在下方的 Properties（属性） 窗格中，寻找 “Memory Pool” 或 “Memory Type” 字段。
如果显示为 Pageable，那么你即便用了 Async 接口，驱动程序依然会在后台进行一次从“用户内存”到“驱动内部锁页缓冲区”的额外同步拷贝。

方法 C：观察传输速率（经验法）

Pinned Memory： 在 PCIe Gen3 x16 上通常能跑到 10-12 GB/s 以上。
Pageable Memory： 往往只有 2-5 GB/s，且波动剧烈。如果你的传输带宽远低于理论值，大概率没用 Pin memory。

为什么 `memcpy` 占比不高，但 Gap 依然很大？

这是一个很隐蔽的逻辑：你看到的 memcpy 耗时是 GPU 引擎搬数的时间，而 Gap 损失的是 CPU 调度的时间。

如果你使用的是 Pageable Memory，cudaMemcpyAsync 的执行流程如下：

CPU 调用： CPU 执行到该行代码。
驱动同步操作： 因为内存是可分页的，CUDA 驱动不敢直接让 DMA 引擎去搬（万一搬一半内存被操作系统移走了怎么办？）。
隐形阻塞： 驱动会迫使 CPU 等待，先将数据拷贝到一块内部的临时 Pinned Buffer 中。
函数返回： 只有等这个内部拷贝完成了，cudaMemcpyAsync 才会返回。

结果就是： 虽然在 GPU 轴上 memcpy 看起来很快，但 CPU 却被卡在了这个 API 调用里，无法及时下发接下来的 Kernel Launch 指令。这在 Trace 上就表现为 Memcpy 结束后，隔了很久才出现 Kernel。

特别注意：ptsz (Per-Thread Default Stream)

cudaMemcpyAsync_ptsz_v7000 揭示了一个重要信息：你的代码编译时可能开启了 CUDA Per-Thread Default Stream 选项（或者框架默认开启）。

优点： 不同的主机线程拥有独立的默认流，减少了全局锁竞争。
潜在风险： 如果这 9-11 个模型之间存在某种依赖，或者它们在不经意间触碰了零号流（Default Stream 0），会引发隐式的全局同步。

观察 stream 0：

Name	Wall duration	Self duration	Avg Wall duration	Occurrences
synchronization(STREAM_SYNCHRONIZE)	59.27ms	58.70ms	29.77µs	1991
synchronization(STREAM_WAIT_EVENT)	12.25ms	12.24ms	1.882µs	6511

有两个较大问题：

STREAM_SYNCHRONIZE (59.27ms / 1991 次)：含义是 CPU 线程在等待 GPU 完成某个 Stream 上的所有任务。平均耗时每次约 $30\mu s$ 。看起来很短，但发生了 近 2000 次；在动态凑批场景下，这意味着 CPU 调度器**每处理一个动作（或一小批请求）就会停下来检查 GPU 是否完成。这正是我们看到 Gap 的直接原因。每当 CPU 调用 cudaStreamSynchronize，GPU 即使计算完了，CPU 也需要时间唤醒、处理业务逻辑、再重新下发 Kernel。这几千次微小的停顿累积成了巨大的空隙。
密集的 STREAM_WAIT_EVENT (6511次)：该函数表示 GPU 内部的同步（Stream A 等待 Stream B 完成某个 Event）。6500 多次调用说明我们的多模型混跑逻辑中，模型之间的依赖关系或资源竞争非常复杂。如果模型之间本应独立，不应该有这么多 Wait Event。这通常发生在频繁的资源重分配或显存复用逻辑中。

总结

现象	可能原因	优化方案
Memcpy 紧跟 Kernel 之间有 Gap	驱动在做 Pageable 内存的预处理	必须改用 `cudaMallocHost` 分配输入 Buffer
Kernel 与 Kernel 之间有 Gap	CPU 凑批逻辑太慢，或 Python/Go 的 GIL/调度损耗	检查 CPU 侧逻辑，考虑用 C++ 重写调度核心
大量 MemsetAsync	频繁初始化显存空间	尽量复用 Buffer，减少 `memset` 次数

针对“多模型混跑”场景的特别意义

在 11 个模型混跑的情况下，如果不使用异步和锁页内存，会出现以下惨状：

模型 A 搬运数据时，CPU 线程被锁死，无法为 模型 B 准备 Batch。
模型 B 好不容易准备好了，发现 模型 A 还在占着 GPU 指令通道下发 Kernel。
所有的模型都在排队等 CPU 释放阻塞，导致 GPU 像在“等红绿灯”一样，走一步停三秒。